[MUSIC]
Hi everyone.
Welcome back to the Heterogeneous Parallel Programming class.
This is lecture 1.6, introduction to CUDA.
We will be discussing the kernel-based SPMD parallel programming.
The objective of this lecture is to help
you to learn the basic concepts involved in a
simple CUDA kernel function, how do we declare a
kernel, how do we use built-in variables in a
kernel, and also to review how we can use map thread index into data index.
Here is a very simple example of how we can declare a CUDA kernel.
And this is the vector addition kernel that
we will eventually be launching from the host code.
But there's several important parts in this kernel.
The first one is the keyword. Every CUDA kernel function
needs to be preceded by __global__ and
this signals to the compiler that the function is a kernel function.
The actual header of the function is exactly
the same as any other traditional C function.
So we have the name in this case is vecAddKernel,
and it's just a name that we gave, we happen
to give to the kernel, and it's followed by all
the parameters that it it is expecting from the caller.
Within the kernel we will be able to use several built-in variables.
The variables that we shown in this in this the code
example are threadIdx.x, blockDim.x and
blockIdx.x. These are all built-in variables that are
initialized by the hardware, so that every thread can access these variables.
And depending on the thread index, and block index, every thread
will have a different combination of the threadIdx and blockIdx values.
So that's why we will be able to assign
a different value to i for every for different threads.
And we already have gone through this concept in a previous lecture.
Finally, we can use
i as a data index to access the the input data.
So in this case, we'll be using i to access one
A element, one B element, and to write into one C element.
We're going to come back and discuss the if statement, why we need to test
what, if i is less than n, n is the number of elements in
the A, B, and C vectors.
This slide shows the kernel launch code in the host code.
In the previous lecture we talk about the allocation of
memory, copying data to the device and then we say
that we are going to to present the kernel launch
part in this lecture and this is where it happens.
So we're going to change the second part of the vecAdd
function in the host code into a kernel launch.
The kernel launch is very similar to any C function calls.
So we still have the C function name, so this is the
vecAddKernel, so this is just the name we assign to the kernel.
But between the function name and the parameters that were
given to the function we also need to supply a configuration
parameter set to the, to the kernel.
So we actually have two configuration parameters.
The first one is the number of blocks in the grid
and the second one is the number of threads in a block.
So what we're saying in this particular statement is that
we want to have 256 threads in each thread block.
And we want to have n divided by 256.0,
and with the ceiling function, number of blocks in the in, in the grid.
So what's happening here is that whenever we want, we generate threads,
we want to have enough threads to cover all the vector elements.
So if we have
256 vector elements, then it's a very simple case.
Because it's 256 threads per block, and we only
need to have one block to cover all the 256.
So we will have 1 and 256 as our configuration parameter.
So, if n is equal to 256, we can simply divide it by 256 and we'll get value 1.
If n
is 512, it's also simple, because 512 divided
by 256 is 2. So, 2 thread blocks will be perfect for
covering 512 elements. The case, that is much more
complicated is when, is, when n is not a multiple of 256.
Let's say if n is 1000. When we divide 1000
by 256 we get 3 point some fraction. And this gives
us a situation where we don't have a perfect multiple of 256.
So in order to, to have enough threads to cover 1000 elements, we actually
need to generate four thread blocks. And so that the
1024 threads altogether will give sufficient coverage to 1000 elements.
And, but then, we will end up with 24 threads that are extra.
So in order to always generate enough threads,
we will divide n by 256.0, and we generate
a floating point number, which is 3 point
a fraction, and then we would take the ceiling
function to to be able to round
up to 4. So this ceiling function ensures that we
always generate the next bigger integer of blocks
that will have enough threads to cover all the elements in this vector.
The reason why we use 256.0 is because we need to make sure that we generate
a floating point number, and that floating number will be rounded up.
If we didn't
use 256.0, but rather if we use 256, the division will be done as an integer
division, so 1000 divided by 256 as an
integer division, will actually generate a value 3.
And then the ceiling function would not be able to take it up to 4.
So this is one of the the ways to
ensure that we can have, always have enough threads.
Now going back to the previous slide.
When we launch, let's say, the kernel to
process 1000 elements, we will have 4 thread blocks.
And we will have a, the last thread block will have
threads that will cover up to a element that element 999, which is the
1000th element because the element start from 0.
However, we will
have 24 threads, where the i value is going to be greater than or equal to 1000,
and these threads really do not correspond to any real elements of A, B and C.
So we don't want those elements to do anything and in fact
if they try to access memory, they can actually corrupt the memory contents.
So, what we want to do is, we want to be able to test whether the i
value for every, for the thread, is
actually less than or equal, less than 1000.
Only threads that succeed in this test should
do the vector, should do the element addition.
So this is a reason why we need to have this check.
And this is commonly referred to as boundary check.
And it's actually testing whether a thread
is supposed to be processing an element within
the boundary or outside the boundary.
So whenever the thread corresponds to an element that is outside the
boundary, then the thread will not be doing anything in the computation.
This slide shows two additional
concepts for the kernel launch. The first one is, we
actually can have two dimensional and three dimensional grids and blocks.
So the particular method that we showed in the previous
slide only works for one dimensional, launching one dimensional grids.
So each of those parameters are actually a integer value.
Here, we show that a, a general method for launching a
grid of 2 dimension, 3 dimension, or in some cases, 1 dimension.
So we, now we need to declare true dimensional variables.
They're all typed dim3, and dim3 is available in the CUDA header file.
And dim3 only means that a variable has three parts.
Each part is a integer value. And we can initialize the variable with
the syntax that we have here by having the variable name followed by parenthesis.
So the first value will initialize the x value, the second value will
initialize the y value, and the third value will initialize the z value.
And in this case, because we're using a 1-dimensional grid,
we only need to initialize the x-values of both the DimGrid and DimBlock.
These variable names are of your choice.
You can call them anything you want, as long as
you use them in the right position of the kernel launch.
So in this case, we call DimGrid the configuration that
we want for the number of blocks in a grid.
And then we
have DimBlock as the number of threads in in the thread block.
So we can use this as a more general way of launching kernels.
The second one is, I'm actually giving you an alternative
way to generate enough thread blocks to cover all the elements.
Here you can use n minus one divided by 256, in this case 256
can be a integer. So this will give us a integer
value, and then you'll plus 1, and this will always give us
the next bigger integer that will cover the, all the threads.
So if we go back to the 1000 example, you know, 1000
divided by 1 will be 999, subtracted by 1 will
be 999, and if you divide it by
256, you will get 3. And 3 plus 1 is 4.
And the most important case is that if you have, let's say, 256, what happens
is, 256 minus 1 is 255, so you will get a 0 from that division, but we add
1. So everything from 0 to 256
will give us exactly one block. In the according to this division.
So if you work out all these possibilities, you'll
realize that they all, this expression always gives you
the exact number the correct number of thread blocks
that will be big enough to cover all the elements.
So now let's put everything together. The kernel execution works as follows.
In the host code, you can, you have a
host function, and then the host function can launch
a kernel by calling the kernel function just like
any other C function, but with additional configuration parameters.
The configuration parameters need to be preceded by three
less signs and followed by three greater signs.
And so once we launched a kernel, we call the kernel function, then the
hardware will generate a grid of threads to execute the kernel function.
Every thread will be executing this kernel function.
And the kernel function is declared in
the CUDA program by preceding the kernel function
decoration with a key word __global__.
So, this allows the host code and the device code to connect, and so that
whenever the host code launch this kernel, it will generate this grid of threads.
And all the threads in the grid will
have these built in variables, blockIdx, blockDim.x, and
threadIdx.x. And these predefined, preinitialized
variables will allow the threads to generate different data indices,
so that each thread will be processing a different part of the code, of the data.
So now, I like to, to just kind of summarize how
you can treat all the function declarations in your CUDA program.
We already talk
about the kernel function and the host function.
So there are really three types of functions in a CUDA program.
The first type is what we have not talked about so far, which is device function.
It's __device__ and this keyword tells the compiler
that the function is going to be a
device function that can only, that can be executed on a device, and it can be called
by another function that is executing on the device.
This kind of function is typically called by
a kernel function or by other device functions.
The second type is what we talk about the most so far, the kernel function.
So the kernel function is a very special device function.
It can be executed on the device, but it can only be called by the host.
This is really the gateway of parallel execution.
When the, when the host code wants to trigger any kind of parallel
execution, it launches a kernel function. And the third type is the host function.
It can be optionally preceded by __host__.
And these kind of functions are just the traditional C functions that will
be executed on the host, and can only be called from the host.
So, a few
important details. __global__ declares
that kernel function. And make sure that you remember the,
there are two underscores before and two underscores after.
And the kernel function must return a void type, because these
functions are special. These functions are really used to
trigger parallel execution on the device.
There's not an easy, simple way to to return a value.
The question is, which thread would even give you that value back.
So that's why to make it simple, the kernel functions must always return void.
And the second one is, the device,
__device__ and __host__ can be used together and
this gives you a function that will be compiled in both ways.
It will be compiled into a host function,
it will be also compiled into a device function.
This way the function can be appropriately used in both the host and device.
And this is oftentimes used, you know used to generate useful
utility functions that could be used on both device and host.
Finally, if the host, if a function is truly only
going to be a host function, then __host__ is really optional.
And this allows you to port a C program
into CUDA very quickly, because all the original programs
can just remain as, with no keywords, and they
will be assumed to run only on the host.
And then you can gradually add keywords
to the function declarations to make kernel functions and device functions
or make some of the functions both host and device functions.
Now that we know how to declare a kernel, how to launch a kernel and so on, I
would like to talk, tell you a little bit more about how a CUDA program is compiled.
In every file, we're going to have both host
function and device function, so the NVCC, which is
the Nvidia C compiler, will automatically identify all the
host functions and kernel functions and separate them out
into two compilation paths.
The host functions will go through the host
code path, where they will be conferred or compiled
by a host C compiler and linker, such as GCC, such as Intel ICC, and so on.
And the second one path is for the device function or the kernel functions.
And these functions are compiled into a binary file format called PTX.
Now, this is very similar
to Java bytecode.
At run time, there is a just-in-time compiler that will take the PTX binary and
generate the real ISA, instruction set architecture, binary
for that particular device that you are using.
So then, the, both types of executables will be used
for this Heterogeneous Computing Platform, and the host path will be
used on the host device which is the CPU, and then the device
binary will be run on the device you know, hardware, such as a GPU.
And this allows you to take a single piece of program with
host and device functions and run
them correctly in a heterogeneous computing platform.
So at this point, we have finished the basic concepts
of, involved in writing a simple, basic CUDA program.
And so I would like to encourage you to begin to take the lab tour,
and the lab tour will show you the the environment
for you to write the CUDA kernels and so on
for the for all the assignments.
And then we would also show you how to get
the instructions and how to submit the your, your assignment.
Once you finish that tour, I like you to start encourage you to start, the
first lab assignment which is the vector addition CUDA code.
So this will help you to solidify your understanding of all the concepts
that we discussed to date.
For those of you who would like to learn more about the CUDA language at the
introductory level, I would like you to, recommend
that you read chapter 3 of the textbook.
Thank you.